//
// Created by xuxu on 2020/7/30.
//
#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <ctime>
#include <iostream>
#include <string>
#include <vector>

#include <sys/time.h>

#include "opencl_runtime.h"

std::string LayoutNames[]
        {
                "NHWC4",
                "NC4HW4",
                "NHC4W4",
        };

size_t SizeOf(DataType data_type)
{
    return data_type == DataType::FLOAT16 ? 2 : 4;
}

long getTimeInUs()
{
    long time;
    struct timeval tv;
    gettimeofday(&tv, nullptr);
    time = tv.tv_sec * 1000000 + tv.tv_usec;
    return time;
}

int GetBiggestDivider(int number, int max_divider)
{
    for (int i = max_divider; i != 0; i--)
        if (number % i == 0)
            return i;
    return 1;
}

float __gnu_h2f_ieee(short param)
{
    unsigned short expHalf16 = param & 0x7C00;
    int exp1 = (int) expHalf16;
    unsigned short mantissa16 = param & 0x03FF;
    int mantissa1 = (int) mantissa16;
    int sign = (int) (param & 0x8000);
    sign = sign << 16;

    // nan or inf
    if (expHalf16 == 0x7C00)
    {
        // nan
        if (mantissa16 > 0)
        {
            int res = (0x7FC00000 | sign);
            float fres = *((float *) (&res));
            return fres;
        }
        // inf
        int res = (0x7F800000 | sign);
        float fres = *((float *) (&res));
        return fres;
    }
    if (expHalf16 != 0)
    {
        exp1 += ((127 - 15) << 10); //exponents converted to float32 bias
        int res = (exp1 | mantissa1);
        res = res << 13;
        res = (res | sign);
        float fres = *((float *) (&res));
        return fres;
    }

    int xmm1 = exp1 > (1 << 10) ? exp1 : (1 << 10);
    xmm1 = (xmm1 << 13);
    xmm1 += ((127 - 15 - 10) << 23);  // add the bias difference to xmm1
    xmm1 = xmm1 | sign; // Combine with the sign mask

    float res = (float) mantissa1;  // Convert mantissa to float
    res *= *((float *) (&xmm1));

    return res;
}

short __gnu_f2h_ieee(float param)
{
    unsigned int param_bit = *((unsigned int *) (&param));
    int sign = param_bit >> 31;
    int mantissa = param_bit & 0x007FFFFF;
    int exp = ((param_bit & 0x7F800000) >> 23) + 15 - 127;
    short res;
    if (exp > 0 && exp < 30)
    {
        // use rte rounding mode, round the significand, combine sign, exponent and significand into a short.
        res = (sign << 15) | (exp << 10) | ((mantissa + 0x00001000) >> 13);
    }
    else if (param_bit == 0)
    {
        res = 0;
    }
    else
    {
        if (exp <= 0)
        {
            if (exp < -10)
            {
                // value is less than min half float point
                res = 0;
            }
            else
            {
                // normalized single, magnitude is less than min normal half float point.
                mantissa = (mantissa | 0x00800000) >> (1 - exp);
                // round to nearest
                if ((mantissa & 0x00001000) > 0)
                {
                    mantissa = mantissa + 0x00002000;
                }
                // combine sign & mantissa (exp is zero to get denormalized number)
                res = (sign << 15) | (mantissa >> 13);
            }
        }
        else if (exp == (255 - 127 + 15))
        {
            if (mantissa == 0)
            {
                // input float is infinity, return infinity half
                res = (sign << 15) | 0x7C00;
            }
            else
            {
                // input float is NaN, return half NaN
                res = (sign << 15) | 0x7C00 | (mantissa >> 13);
            }
        }
        else
        {
            // exp > 0, normalized single, round to nearest
            if ((mantissa & 0x00001000) > 0)
            {
                mantissa = mantissa + 0x00002000;
                if ((mantissa & 0x00800000) > 0)
                {
                    mantissa = 0;
                    exp = exp + 1;
                }
            }
            if (exp > 30)
            {
                // exponent overflow - return infinity half
                res = (sign << 15) | 0x7C00;
            }
            else
            {
                // combine sign, exp and mantissa into normalized half
                res = (sign << 15) | (exp << 10) | (mantissa >> 13);
            }
        }
    }
    return res;
}

std::vector<short> FP32ToFP16(const std::vector<float> &X)
{
    std::vector<short> Y;
    Y.reserve(X.size());
    for (auto x: X)
    {
        Y.push_back(__gnu_f2h_ieee(x));
    }
    return Y;
}

std::vector<float> FP16ToFP32(const std::vector<short> &X)
{
    std::vector<float> Y;
    Y.reserve(X.size());
    for (auto x: X)
    {
        Y.push_back(__gnu_h2f_ieee(x));
    }
    return Y;
}

void printf_callback(const char *buffer, size_t length, size_t final, void *user_data)
{
    fwrite(buffer, 1, length, stdout);
};

std::string ReadKernelSource(const std::string &clFileName)
{
    // .cl string: 读入.cl到字符串 (https://stackoverflow.com/questions/2912520/read-file-contents-into-a-string-in-c)
    FILE *clFile = fopen(clFileName.c_str(), "r");
    fseek(clFile, 0, SEEK_END);
    size_t sourceSize = ftell(clFile);
    std::string kernelSourceCode(sourceSize, 0);
    rewind(clFile);
    fread(&kernelSourceCode[0], sizeof(char), sourceSize, clFile);
    return kernelSourceCode;
}

//int OpenCLRuntime::RunHello()
//{
//    cl_int status;
//    constexpr char clFileName[] = "cl/hello.cl";
//    constexpr char kernelName[] = "hello_cl";
//    BuildKernel(clFileName, kernelName);
//
//    // 创建 OpenCL buffer 对象
//    cl_mem outputBuffer = mslite::clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR, 4 * 4 * 4, nullptr, &status);
//    RETURN_IF_ERROR(status != CL_SUCCESS, "Error: Create Buffer, outputBuffer (clCreateBuffer)\n");
//
//    // 为 kernel 设置参数
//    status = mslite::clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &outputBuffer);
//    RETURN_IF_ERROR(status != CL_SUCCESS, "Error: Setting kernel argument (clSetKernelArg)\n");
//
//    // 将一个 kernel 放入 command queue
//    printf("\ncall clEnqueueNDRangeKernel():\n");
//    size_t globalThreads[] = {4, 4};
//    size_t localThreads[] = {2, 2};
//    status = mslite::clEnqueueNDRangeKernel(commandQueue, kernel, sizeof(globalThreads) / sizeof(size_t), nullptr,
//                                            globalThreads, localThreads, 0, nullptr, nullptr);
//    RETURN_IF_ERROR(status != CL_SUCCESS, "Error: Enqueueing kernel\n");
//
//    // 确认 command queue 中所有命令都执行完毕
//    status = mslite::clFinish(commandQueue);
//    RETURN_IF_ERROR(status != CL_SUCCESS, "Error: Finish command queue\n")
//
//    // 将内存对象中的结果读回 Host
//    std::vector<unsigned int> outbuffer(4 * 4, 0);
//    status = mslite::clEnqueueReadBuffer(commandQueue, outputBuffer, CL_TRUE, 0, 4 * 4 * 4, outbuffer.data(), 0,
//                                         nullptr, nullptr);
//    RETURN_IF_ERROR(status != CL_SUCCESS, "Error: Read buffer queue\n")
//
//    // Host 端打印结果
//    // 预期结果
//    // 0x11 0x12 0x14 0x18
//    // 0x21 0x22 0x24 0x28
//    // 0x41 0x42 0x44 0x48
//    // 0x81 0x82 0x84 0x88
//    printf("\noutput:\n");
//    for (int i = 0; i < 16; ++i)
//    {
//        printf("0x%x ", outbuffer[i]);
//        if ((i + 1) % 4 == 0)
//            printf("\n");
//    }
//    status = mslite::clReleaseMemObject(outputBuffer);
//
//    return 0;
//}